
#include <stdio.h>
#include <stdlib.h>

#define N  1048576
#define THREADS_PER_BLOCK 512

// 交错配对的规约求和核函数,使用全局内存
__global__ void reduceInterleaved(float *g_idata, float *g_odata, int n) {
    unsigned int tid = threadIdx.x;
    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 每个线程处理一个数据元素,idata指向当前线程块要处理的第一个数据元素的位置
    float *idata = g_idata + blockIdx.x * blockDim.x;
    
    // 规约操作
    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
        if (tid < stride && idx + stride < n) {
            idata[tid] += idata[tid + stride];
        }
        __syncthreads();
    }
    
    // 将结果写入输出数组
    if (tid == 0) {
        g_odata[blockIdx.x] = idata[0];
    }
}

int main() {
    float *h_array = (float *)malloc(N * sizeof(float));
    float *d_array, *d_result;
    float final_result;
    
    // 初始化输入数据
    for (int i = 0; i < N; i++) {
        h_array[i] = 1.0f;
    }
    
    // 分配设备内存
    cudaMalloc(&d_array, N * sizeof(float));
    cudaMalloc(&d_result, (N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK * sizeof(float));
    
    // 拷贝数据到设备
    cudaMemcpy(d_array, h_array, N * sizeof(float), cudaMemcpyHostToDevice);
    
    // 第一次规约调用
    dim3 blocks((N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK);
    dim3 threads(THREADS_PER_BLOCK);
    reduceInterleaved<<<blocks, threads>>>(d_array, d_result, N);

    // 第二次规约调用(处理第一次的结果)
    int newsize = blocks.x;
    if(newsize <= THREADS_PER_BLOCK) {
        reduceInterleaved<<<1, newsize>>>(d_result, d_result, newsize);
    } else {
        // 如果newsize太大，仍需要循环处理
        while(newsize > 1) {
            dim3 newBlocks((newsize + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK);
            reduceInterleaved<<<newBlocks, threads>>>(d_result, d_result, newsize);
            newsize = newBlocks.x;
        }
    }
      
    // 拷贝最终结果回主机
    cudaMemcpy(&final_result, d_result, sizeof(float), cudaMemcpyDeviceToHost);
    
    printf("Final sum: %f\n", final_result);
    
    // 释放内存
    free(h_array);
    cudaFree(d_array);
    cudaFree(d_result);
    
    return 0;
}
